// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// Copyright (C) 2014 Benoit Steiner <benoit.steiner.goog@gmail.com>
//
// This Source Code Form is subject to the terms of the Mozilla
// Public License v. 2.0. If a copy of the MPL was not distributed
// with this file, You can obtain one at http://mozilla.org/MPL/2.0/.

#ifndef EIGEN_CXX11_TENSOR_TENSOR_EXECUTOR_H
#define EIGEN_CXX11_TENSOR_TENSOR_EXECUTOR_H

namespace Eigen {

/**
 * \class TensorExecutor
 * \ingroup CXX11_Tensor_Module
 *
 * \brief The tensor executor class.
 *
 * This class is responsible for launch the evaluation of the expression on
 * the specified computing device.
 *
 * @tparam Vectorizable can use packet math (SSE/AVX/etc... registers and
 *                      instructions)
 * @tparam Tiling       can use block based tensor evaluation
 *                      (see TensorBlock.h)
 */
namespace internal {

/**
 * Evaluating TensorBroadcastingOp via coefficient of packet path is extremely
 * expensive. If expression has at least one broadcast op in it, and it supports
 * block based evaluation, we always prefer it, even for the small tensors. For
 * all other tileable ops, block evaluation overhead for small tensors (fits
 * into L1) is too large, and we fallback on vectorized evaluation.
 */

// TODO(ezhulenev): Add specializations for all other types of Tensor ops.

template<typename Expression>
struct ExpressionHasTensorBroadcastingOp
{
	enum
	{
		value = false
	};
};

template<typename LhsXprType, typename RhsXprType>
struct ExpressionHasTensorBroadcastingOp<const TensorAssignOp<LhsXprType, RhsXprType>>
{
	enum
	{
		value = ExpressionHasTensorBroadcastingOp<RhsXprType>::value
	};
};

template<typename UnaryOp, typename XprType>
struct ExpressionHasTensorBroadcastingOp<const TensorCwiseUnaryOp<UnaryOp, XprType>>
{
	enum
	{
		value = ExpressionHasTensorBroadcastingOp<XprType>::value
	};
};

template<typename BinaryOp, typename LhsXprType, typename RhsXprType>
struct ExpressionHasTensorBroadcastingOp<const TensorCwiseBinaryOp<BinaryOp, LhsXprType, RhsXprType>>
{
	enum
	{
		value =
			ExpressionHasTensorBroadcastingOp<LhsXprType>::value || ExpressionHasTensorBroadcastingOp<RhsXprType>::value
	};
};

template<typename Broadcast, typename XprType>
struct ExpressionHasTensorBroadcastingOp<const TensorBroadcastingOp<Broadcast, XprType>>
{
	enum
	{
		value = true
	};
};

// -------------------------------------------------------------------------- //

/**
 * Default strategy: the expression is evaluated sequentially with a single cpu
 * thread, without vectorization and block evaluation.
 */
template<typename Expression, typename Device, bool Vectorizable, TiledEvaluation Tiling>
class TensorExecutor
{
  public:
	typedef typename Expression::Index StorageIndex;

	// Including `unsupported/Eigen/CXX11/Tensor` in different translation units
	// with/without `EIGEN_USE_THREADS` or `EIGEN_USE_GPU` is a potential ODR
	// violation. If this template is instantiated with a non-default device, it
	// means that this header file was included without defining
	// `EIGEN_USE_THREADS`, `EIGEN_USE_GPU` or `EIGEN_USE_SYCL`.
	static_assert(std::is_same<Device, DefaultDevice>::value,
				  "Default executor instantiated with non-default device. "
				  "You must #define EIGEN_USE_THREADS, EIGEN_USE_GPU or "
				  "EIGEN_USE_SYCL before including Eigen headers.");

	EIGEN_DEVICE_FUNC
	static EIGEN_STRONG_INLINE void run(const Expression& expr, const Device& device = Device())
	{
		TensorEvaluator<Expression, Device> evaluator(expr, device);
		const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL);
		if (needs_assign) {
			const StorageIndex size = array_prod(evaluator.dimensions());
			for (StorageIndex i = 0; i < size; ++i) {
				evaluator.evalScalar(i);
			}
		}
		evaluator.cleanup();
	}
};

/**
 * Default async execution strategy is not implemented. Currently it's only
 * available for ThreadPoolDevice (see definition below).
 */
template<typename Expression, typename Device, typename DoneCallback, bool Vectorizable, TiledEvaluation Tiling>
class TensorAsyncExecutor
{};

/**
 * Process all the data with a single cpu thread, using vectorized instructions.
 */
template<typename Expression>
class TensorExecutor<Expression,
					 DefaultDevice,
					 /*Vectorizable=*/true,
					 /*Tiling=*/TiledEvaluation::Off>
{
  public:
	typedef typename Expression::Index StorageIndex;

	EIGEN_DEVICE_FUNC
	static EIGEN_STRONG_INLINE void run(const Expression& expr, const DefaultDevice& device = DefaultDevice())
	{
		TensorEvaluator<Expression, DefaultDevice> evaluator(expr, device);
		const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL);
		if (needs_assign) {
			const StorageIndex size = array_prod(evaluator.dimensions());
			const int PacketSize =
				unpacket_traits<typename TensorEvaluator<Expression, DefaultDevice>::PacketReturnType>::size;

			// Give compiler a strong possibility to unroll the loop. But don't insist
			// on unrolling, because if the function is expensive compiler should not
			// unroll the loop at the expense of inlining.
			const StorageIndex UnrolledSize = (size / (4 * PacketSize)) * 4 * PacketSize;
			for (StorageIndex i = 0; i < UnrolledSize; i += 4 * PacketSize) {
				for (StorageIndex j = 0; j < 4; j++) {
					evaluator.evalPacket(i + j * PacketSize);
				}
			}
			const StorageIndex VectorizedSize = (size / PacketSize) * PacketSize;
			for (StorageIndex i = UnrolledSize; i < VectorizedSize; i += PacketSize) {
				evaluator.evalPacket(i);
			}
			for (StorageIndex i = VectorizedSize; i < size; ++i) {
				evaluator.evalScalar(i);
			}
		}
		evaluator.cleanup();
	}
};

/**
 * Process all the data with a single cpu thread, using blocks of data. By
 * sizing a block to fit L1 cache we get better cache performance.
 */
template<typename Expression, bool Vectorizable>
class TensorExecutor<Expression,
					 DefaultDevice,
					 Vectorizable,
					 /*Tiling=*/TiledEvaluation::On>
{
  public:
	typedef typename traits<Expression>::Scalar Scalar;
	typedef typename remove_const<Scalar>::type ScalarNoConst;

	typedef TensorEvaluator<Expression, DefaultDevice> Evaluator;
	typedef typename traits<Expression>::Index StorageIndex;

	static const int NumDims = traits<Expression>::NumDimensions;

	EIGEN_DEVICE_FUNC
	static EIGEN_STRONG_INLINE void run(const Expression& expr, const DefaultDevice& device = DefaultDevice())
	{
		typedef TensorBlockMapper<NumDims, Evaluator::Layout, StorageIndex> TensorBlockMapper;

		typedef internal::TensorBlockDescriptor<NumDims, StorageIndex> TensorBlockDesc;
		typedef internal::TensorBlockScratchAllocator<DefaultDevice> TensorBlockScratch;

		Evaluator evaluator(expr, device);

		// TODO(ezhulenev): Do not use tiling for small tensors?
		const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL);

		if (needs_assign) {
			// Query expression tree for desired block size/shape.
			const TensorBlockResourceRequirements requirements = evaluator.getResourceRequirements();

			const TensorBlockMapper block_mapper(typename TensorBlockDesc::Dimensions(evaluator.dimensions()),
												 requirements);

			// Share scratch memory allocator between all blocks.
			TensorBlockScratch scratch(device);

			const StorageIndex total_block_count = block_mapper.blockCount();
			for (StorageIndex i = 0; i < total_block_count; ++i) {
				TensorBlockDesc desc = block_mapper.blockDescriptor(i);
				evaluator.evalBlock(desc, scratch);
				scratch.reset();
			}
		}
		evaluator.cleanup();
	}
};

/**
 * Multicore strategy: the index space is partitioned and each partition is
 * executed on a single core.
 *
 * (1) TensorExecutor will submit work to the ThreadPoolDevice managed thread
 *     pool, and will block the caller thread until all tasks are finished.
 *
 * (2) TensorAsyncExecutor is a non-blocking version, that will submit work to
 *     the ThreadPoolDevice managed thread pool, and will return immediately.
 *     It will call 'done' callback after all tasks are finished.
 */
#ifdef EIGEN_USE_THREADS

template<typename TensorBlockMapper>
struct TensorExecutorTilingContext
{
	TensorExecutorTilingContext() = default;
	TensorExecutorTilingContext(const TensorBlockMapper& b_mapper, const TensorOpCost& b_cost, size_t b_aligned_size)
		: block_mapper(b_mapper)
		, cost(b_cost)
		, aligned_blocksize(b_aligned_size)
	{
	}

	TensorBlockMapper block_mapper; // navigate through blocks
	TensorOpCost cost;				// cost of computing a single block
	size_t aligned_blocksize;		// block size after memory alignment
};

// Computes a block evaluation parameters, and allocates temporary memory buffer
// for blocks. See TensorExecutor/TensorAsyncExecutor (Tiling=On) below.
template<typename Evaluator, typename TensorBlockMapper, bool Vectorizable>
TensorExecutorTilingContext<TensorBlockMapper>
GetTensorExecutorTilingContext(const Evaluator& evaluator)
{
	// Query expression tree for desired block size/shape.
	TensorBlockResourceRequirements requirements = evaluator.getResourceRequirements();

	// Update target block size based on cost model.
	double taskSize = TensorCostModel<ThreadPoolDevice>::taskSize(1, requirements.cost_per_coeff);
	requirements.size = static_cast<size_t>(1.0 / taskSize);

	TensorBlockMapper block_mapper(typename TensorBlockMapper::Dimensions(evaluator.dimensions()), requirements);

	size_t block_size = block_mapper.blockTotalSize();
	const size_t align = numext::maxi(EIGEN_MAX_ALIGN_BYTES, 1);
	const size_t aligned_blocksize = align * divup<size_t>(block_size * sizeof(typename Evaluator::Scalar), align);

	return { block_mapper, requirements.cost_per_coeff * block_size, aligned_blocksize };
}

template<typename Evaluator, typename StorageIndex, bool Vectorizable>
struct EvalRange
{
	static void run(Evaluator* evaluator_in, const StorageIndex firstIdx, const StorageIndex lastIdx)
	{
		Evaluator evaluator = *evaluator_in;
		eigen_assert(lastIdx >= firstIdx);
		for (StorageIndex i = firstIdx; i < lastIdx; ++i) {
			evaluator.evalScalar(i);
		}
	}

	static StorageIndex alignBlockSize(StorageIndex size) { return size; }
};

template<typename Evaluator, typename StorageIndex>
struct EvalRange<Evaluator, StorageIndex, /*Vectorizable*/ true>
{
	static const int PacketSize = unpacket_traits<typename Evaluator::PacketReturnType>::size;

	static void run(Evaluator* evaluator_in, const StorageIndex firstIdx, const StorageIndex lastIdx)
	{
		Evaluator evaluator = *evaluator_in;
		eigen_assert(lastIdx >= firstIdx);
		StorageIndex i = firstIdx;
		if (lastIdx - firstIdx >= PacketSize) {
			eigen_assert(firstIdx % PacketSize == 0);
			StorageIndex last_chunk_offset = lastIdx - 4 * PacketSize;
			// Give compiler a strong possibility to unroll the loop. But don't insist
			// on unrolling, because if the function is expensive compiler should not
			// unroll the loop at the expense of inlining.
			for (; i <= last_chunk_offset; i += 4 * PacketSize) {
				for (StorageIndex j = 0; j < 4; j++) {
					evaluator.evalPacket(i + j * PacketSize);
				}
			}
			last_chunk_offset = lastIdx - PacketSize;
			for (; i <= last_chunk_offset; i += PacketSize) {
				evaluator.evalPacket(i);
			}
		}
		for (; i < lastIdx; ++i) {
			evaluator.evalScalar(i);
		}
	}

	static StorageIndex alignBlockSize(StorageIndex size)
	{
		// Align block size to packet size and account for unrolling in run above.
		if (size >= 16 * PacketSize) {
			return (size + 4 * PacketSize - 1) & ~(4 * PacketSize - 1);
		}
		// Aligning to 4 * PacketSize would increase block size by more than 25%.
		return (size + PacketSize - 1) & ~(PacketSize - 1);
	}
};

template<typename Expression, bool Vectorizable, TiledEvaluation Tiling>
class TensorExecutor<Expression, ThreadPoolDevice, Vectorizable, Tiling>
{
  public:
	typedef typename Expression::Index StorageIndex;

	static EIGEN_STRONG_INLINE void run(const Expression& expr, const ThreadPoolDevice& device)
	{
		typedef TensorEvaluator<Expression, ThreadPoolDevice> Evaluator;
		typedef EvalRange<Evaluator, StorageIndex, Vectorizable> EvalRange;

		Evaluator evaluator(expr, device);
		const bool needs_assign = evaluator.evalSubExprsIfNeeded(nullptr);
		if (needs_assign) {
			const StorageIndex size = array_prod(evaluator.dimensions());
			device.parallelFor(size,
							   evaluator.costPerCoeff(Vectorizable),
							   EvalRange::alignBlockSize,
							   [&evaluator](StorageIndex firstIdx, StorageIndex lastIdx) {
								   EvalRange::run(&evaluator, firstIdx, lastIdx);
							   });
		}
		evaluator.cleanup();
	}
};

template<typename Expression, bool Vectorizable>
class TensorExecutor<Expression,
					 ThreadPoolDevice,
					 Vectorizable,
					 /*Tiling=*/TiledEvaluation::On>
{
  public:
	typedef typename traits<Expression>::Index IndexType;
	typedef typename traits<Expression>::Scalar Scalar;
	typedef typename remove_const<Scalar>::type ScalarNoConst;

	static const int NumDims = traits<Expression>::NumDimensions;

	typedef TensorEvaluator<Expression, ThreadPoolDevice> Evaluator;
	typedef TensorBlockMapper<NumDims, Evaluator::Layout, IndexType> BlockMapper;
	typedef TensorExecutorTilingContext<BlockMapper> TilingContext;

	typedef internal::TensorBlockDescriptor<NumDims, IndexType> TensorBlockDesc;
	typedef internal::TensorBlockScratchAllocator<ThreadPoolDevice> TensorBlockScratch;

	static EIGEN_STRONG_INLINE void run(const Expression& expr, const ThreadPoolDevice& device)
	{
		Evaluator evaluator(expr, device);

		const bool needs_assign = evaluator.evalSubExprsIfNeeded(nullptr);
		if (needs_assign) {
			const TilingContext tiling =
				internal::GetTensorExecutorTilingContext<Evaluator, BlockMapper, Vectorizable>(evaluator);

			auto eval_block = [&device, &evaluator, &tiling](IndexType firstBlockIdx, IndexType lastBlockIdx) {
				TensorBlockScratch scratch(device);

				for (IndexType block_idx = firstBlockIdx; block_idx < lastBlockIdx; ++block_idx) {
					TensorBlockDesc desc = tiling.block_mapper.blockDescriptor(block_idx);
					evaluator.evalBlock(desc, scratch);
					scratch.reset();
				}
			};

			// Evaluate small expressions directly as a single block.
			if (tiling.block_mapper.blockCount() == 1) {
				TensorBlockScratch scratch(device);
				TensorBlockDesc desc(0, tiling.block_mapper.blockDimensions());
				evaluator.evalBlock(desc, scratch);
			} else {
				device.parallelFor(tiling.block_mapper.blockCount(), tiling.cost, eval_block);
			}
		}
		evaluator.cleanup();
	}
};

template<typename Expression, typename DoneCallback, bool Vectorizable, TiledEvaluation Tiling>
class TensorAsyncExecutor<Expression, ThreadPoolDevice, DoneCallback, Vectorizable, Tiling>
{
  public:
	typedef typename Expression::Index StorageIndex;
	typedef TensorEvaluator<Expression, ThreadPoolDevice> Evaluator;

	static EIGEN_STRONG_INLINE void runAsync(const Expression& expr, const ThreadPoolDevice& device, DoneCallback done)
	{
		TensorAsyncExecutorContext* const ctx = new TensorAsyncExecutorContext(expr, device, std::move(done));

		const auto on_eval_subexprs = [ctx, &device](bool need_assign) -> void {
			if (!need_assign) {
				delete ctx;
				return;
			}

			typedef EvalRange<Evaluator, StorageIndex, Vectorizable> EvalRange;
			const StorageIndex size = array_prod(ctx->evaluator.dimensions());
			device.parallelForAsync(
				size,
				ctx->evaluator.costPerCoeff(Vectorizable),
				EvalRange::alignBlockSize,
				[ctx](StorageIndex firstIdx, StorageIndex lastIdx) {
					EvalRange::run(&ctx->evaluator, firstIdx, lastIdx);
				},
				[ctx]() { delete ctx; });
		};

		ctx->evaluator.evalSubExprsIfNeededAsync(nullptr, on_eval_subexprs);
	}

  private:
	struct TensorAsyncExecutorContext
	{
		TensorAsyncExecutorContext(const Expression& expr, const ThreadPoolDevice& thread_pool, DoneCallback done)
			: evaluator(expr, thread_pool)
			, on_done(std::move(done))
		{
		}

		~TensorAsyncExecutorContext()
		{
			evaluator.cleanup();
			on_done();
		}

		Evaluator evaluator;

	  private:
		DoneCallback on_done;
	};
};

template<typename Expression, typename DoneCallback, bool Vectorizable>
class TensorAsyncExecutor<Expression, ThreadPoolDevice, DoneCallback, Vectorizable, /*Tileable*/ TiledEvaluation::On>
{
  public:
	typedef typename traits<Expression>::Index IndexType;
	typedef typename traits<Expression>::Scalar Scalar;
	typedef typename remove_const<Scalar>::type ScalarNoConst;

	static const int NumDims = traits<Expression>::NumDimensions;

	typedef TensorEvaluator<Expression, ThreadPoolDevice> Evaluator;
	typedef TensorBlockMapper<NumDims, Evaluator::Layout, IndexType> BlockMapper;
	typedef TensorExecutorTilingContext<BlockMapper> TilingContext;

	typedef internal::TensorBlockDescriptor<NumDims, IndexType> TensorBlockDesc;
	typedef internal::TensorBlockScratchAllocator<ThreadPoolDevice> TensorBlockScratch;

	static EIGEN_STRONG_INLINE void runAsync(const Expression& expr, const ThreadPoolDevice& device, DoneCallback done)
	{

		TensorAsyncExecutorContext* const ctx = new TensorAsyncExecutorContext(expr, device, std::move(done));

		const auto on_eval_subexprs = [ctx](bool need_assign) -> void {
			if (!need_assign) {
				delete ctx;
				return;
			}

			ctx->tiling =
				internal::GetTensorExecutorTilingContext<Evaluator, BlockMapper, Vectorizable>(ctx->evaluator);

			auto eval_block = [ctx](IndexType firstBlockIdx, IndexType lastBlockIdx) {
				TensorBlockScratch scratch(ctx->device);

				for (IndexType block_idx = firstBlockIdx; block_idx < lastBlockIdx; ++block_idx) {
					TensorBlockDesc desc = ctx->tiling.block_mapper.blockDescriptor(block_idx);
					ctx->evaluator.evalBlock(desc, scratch);
					scratch.reset();
				}
			};

			// Evaluate small expressions directly as a single block.
			if (ctx->tiling.block_mapper.blockCount() == 1) {
				TensorBlockScratch scratch(ctx->device);
				TensorBlockDesc desc(0, ctx->tiling.block_mapper.blockDimensions());
				ctx->evaluator.evalBlock(desc, scratch);
				delete ctx;
			} else {
				ctx->device.parallelForAsync(
					ctx->tiling.block_mapper.blockCount(), ctx->tiling.cost, eval_block, [ctx]() { delete ctx; });
			}
		};

		ctx->evaluator.evalSubExprsIfNeededAsync(nullptr, on_eval_subexprs);
	}

  private:
	struct TensorAsyncExecutorContext
	{
		TensorAsyncExecutorContext(const Expression& expr, const ThreadPoolDevice& thread_pool, DoneCallback done)
			: device(thread_pool)
			, evaluator(expr, thread_pool)
			, on_done(std::move(done))
		{
		}

		~TensorAsyncExecutorContext()
		{
			evaluator.cleanup();
			on_done();
		}

		const ThreadPoolDevice& device;
		Evaluator evaluator;
		TilingContext tiling;

	  private:
		DoneCallback on_done;
	};
};

#endif // EIGEN_USE_THREADS

// GPU: the evaluation of the expression is offloaded to a GPU.
#if defined(EIGEN_USE_GPU)

template<typename Expression, bool Vectorizable, TiledEvaluation Tiling>
class TensorExecutor<Expression, GpuDevice, Vectorizable, Tiling>
{
  public:
	typedef typename Expression::Index StorageIndex;
	static void run(const Expression& expr, const GpuDevice& device);
};

#if defined(EIGEN_GPUCC)
template<typename Evaluator, typename StorageIndex, bool Vectorizable>
struct EigenMetaKernelEval
{
	static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void run(Evaluator& eval,
														  StorageIndex firstIdx,
														  StorageIndex lastIdx,
														  StorageIndex step_size)
	{
		for (StorageIndex i = firstIdx; i < lastIdx; i += step_size) {
			eval.evalScalar(i);
		}
	}
};

template<typename Evaluator, typename StorageIndex>
struct EigenMetaKernelEval<Evaluator, StorageIndex, true>
{
	static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void run(Evaluator& eval,
														  StorageIndex firstIdx,
														  StorageIndex lastIdx,
														  StorageIndex step_size)
	{
		const StorageIndex PacketSize = unpacket_traits<typename Evaluator::PacketReturnType>::size;
		const StorageIndex vectorized_size = (lastIdx / PacketSize) * PacketSize;
		const StorageIndex vectorized_step_size = step_size * PacketSize;

		// Use the vector path
		for (StorageIndex i = firstIdx * PacketSize; i < vectorized_size; i += vectorized_step_size) {
			eval.evalPacket(i);
		}
		for (StorageIndex i = vectorized_size + firstIdx; i < lastIdx; i += step_size) {
			eval.evalScalar(i);
		}
	}
};

template<typename Evaluator, typename StorageIndex>
__global__ void __launch_bounds__(1024) EigenMetaKernel(Evaluator eval, StorageIndex size)
{

	const StorageIndex first_index = blockIdx.x * blockDim.x + threadIdx.x;
	const StorageIndex step_size = blockDim.x * gridDim.x;

	const bool vectorizable = Evaluator::PacketAccess & Evaluator::IsAligned;
	EigenMetaKernelEval<Evaluator, StorageIndex, vectorizable>::run(eval, first_index, size, step_size);
}

/*static*/
template<typename Expression, bool Vectorizable, TiledEvaluation Tiling>
EIGEN_STRONG_INLINE void
TensorExecutor<Expression, GpuDevice, Vectorizable, Tiling>::run(const Expression& expr, const GpuDevice& device)
{
	TensorEvaluator<Expression, GpuDevice> evaluator(expr, device);
	const bool needs_assign = evaluator.evalSubExprsIfNeeded(nullptr);
	if (needs_assign) {

		const int block_size = device.maxGpuThreadsPerBlock();
		const int max_blocks = device.getNumGpuMultiProcessors() * device.maxGpuThreadsPerMultiProcessor() / block_size;
		const StorageIndex size = array_prod(evaluator.dimensions());
		// Create a least one block to ensure we won't crash when tensorflow calls with tensors of size 0.
		const int num_blocks = numext::maxi<int>(numext::mini<int>(max_blocks, divup<int>(size, block_size)), 1);

		LAUNCH_GPU_KERNEL((EigenMetaKernel<TensorEvaluator<Expression, GpuDevice>, StorageIndex>),
						  num_blocks,
						  block_size,
						  0,
						  device,
						  evaluator,
						  size);
	}
	evaluator.cleanup();
}

#endif // EIGEN_GPUCC
#endif // EIGEN_USE_GPU

// SYCL Executor policy
#ifdef EIGEN_USE_SYCL

template<typename Evaluator>
struct ExecExprFunctorKernel
{
	typedef typename Evaluator::Index Index;
	Evaluator evaluator;
	const Index range;
	template<typename Scratch>
	EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE ExecExprFunctorKernel(const Scratch, Evaluator evaluator_, const Index range_)
		: evaluator(evaluator_)
		, range(range_)
	{
	}

	EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void operator()(cl::sycl::nd_item<1> itemID) { compute(itemID); }
	template<bool is_vec = Evaluator::PacketAccess>
	EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE typename std::enable_if<!is_vec>::type compute(
		const cl::sycl::nd_item<1>& itemID)
	{
		Index gId = static_cast<Index>(itemID.get_global_linear_id());
		Index total_threads = itemID.get_global_range(0);

		for (Index i = gId; i < range; i += total_threads) {
			evaluator.evalScalar(i);
		}
	}
	template<bool is_vec = Evaluator::PacketAccess>
	EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE typename std::enable_if<is_vec>::type compute(
		const cl::sycl::nd_item<1>& itemID)
	{
		const Index vectorizedRange = (range / Evaluator::PacketSize) * Evaluator::PacketSize;
		Index gId = static_cast<Index>(itemID.get_global_linear_id());
		const Index step = Evaluator::PacketSize * itemID.get_global_range(0);
		const Index start = Evaluator::PacketSize * gId;
		for (Index i = start; i < vectorizedRange; i += step) {
			evaluator.evalPacket(i);
		}
		gId += vectorizedRange;
		for (Index i = gId; i < range; i += itemID.get_global_range(0)) {
			evaluator.evalScalar(i);
		}
	}
};

template<typename Expression, bool Vectorizable, TiledEvaluation Tiling>
class TensorExecutor<Expression, Eigen::SyclDevice, Vectorizable, Tiling>
{
  public:
	typedef typename Expression::Index Index;
	static EIGEN_STRONG_INLINE void run(const Expression& expr, const Eigen::SyclDevice& dev)
	{
		typedef Eigen::TensorEvaluator<Expression, Eigen::SyclDevice> Evaluator;
		Evaluator evaluator(expr, dev);
		const bool needs_assign = evaluator.evalSubExprsIfNeeded(NULL);
		if (needs_assign) {
			Index range, GRange, tileSize;
			Index total_size = ::Eigen::internal::array_prod(evaluator.dimensions());
			total_size = (total_size == 0) ? 1 : total_size;
			const int PacketSize = Eigen::PacketType<typename Evaluator::CoeffReturnType, Eigen::SyclDevice>::size;
			Index vectorizable_threads = static_cast<Index>(total_size / PacketSize);
			dev.parallel_for_setup(vectorizable_threads, tileSize, range, GRange);
			range = total_size;

			dev.template nullary_kernel_launcher<typename Evaluator::CoeffReturnType, ExecExprFunctorKernel<Evaluator>>(
				evaluator,
				cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)),
				Index(1),
				range);
		}
		evaluator.cleanup();
	}
};

#endif

} // end namespace internal

} // end namespace Eigen

#endif // EIGEN_CXX11_TENSOR_TENSOR_EXECUTOR_H
